[CK_TILE][FMHA] Add new tile size for async#3586
Conversation
Signed-off-by: Linjun-AMD <Jun.Lin@amd.com>
LJ-underdog
requested review from
Snektron,
ThomasNing,
afagaj,
andriy-ca,
aosewski,
asleepzzz,
bartekxk,
carlushuang,
cgmillette,
coderfeli,
geyyer,
illsilin,
poyenc,
qianfengz,
shumway,
tenpercent,
vidyasagar-amd and
vpietila-amd
as code owners
|
Performance` improve 30% on gfx950 when block_num <= cu_num `./bin/tile_example_fmha_fwd -h=8 -d=128 -s=512 -kname=1 -v=1 ./bin/tile_example_fmha_fwd -h=8 -d=128 -s=512 -kname=1 -v=1 |
There was a problem hiding this comment.
Pull request overview
This PR adds a new tile size configuration for async pipeline operations in the FMHA (Fused Multi-Head Attention) forward pass implementation. The changes introduce support for a 64x128 tile configuration for specific head dimension scenarios and adjust the sequence tuning logic to accommodate this new tile size.
Changes:
- Modified sequence tuning logic to include tile size 64 as a special case alongside the maximum tile size
- Added filtering logic to exclude 64-size tiles for non-async pipelines with 128x128 head dimensions
- Introduced a new 64x128x32 tile size configuration with compute unit constraint for 128x128 head dimensions
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
Signed-off-by: Linjun-AMD <Jun.Lin@amd.com>
This reverts commit f3aafb9.
* add new tile size for async Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> * Update example/ck_tile/01_fmha/codegen/ops/fmha_fwd.py Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> * fix lse error Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> --------- Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
* Revert "Revert "[CK_TILE][FMHA] Add new tile size for async (#3586)" (#3613)" This reverts commit 8f75869. * Add new tile_size for async pipeline Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> * Update include/ck_tile/ops/fmha/pipeline/block_fmha_pipeline_qr_ks_vs_async.hpp Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com> --------- Signed-off-by: Linjun-AMD <Jun.Lin@amd.com> Co-authored-by: Copilot <175728472+Copilot@users.noreply.github.com>
2 participants
Proposed changes
This PR adds a new tile size configuration for async pipeline operations in the FMHA (Fused Multi-Head Attention) forward pass implementation. The changes introduce support for a 64x128 tile configuration for specific head dimension scenarios and adjust the sequence tuning logic to accommodate this new tile size.
Changes:
Modified sequence tuning logic to include tile size 64 as a special case alongside the maximum tile size
Added filtering logic to exclude 64-size tiles for non-async pipelines with 128x128 head dimensions
Introduced a new 64x128x32 tile size configuration with compute unit constraint for 128x128 head dimensions
Checklist
Please put an
xinto the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.clang-formaton all changed filesDiscussion
If this is a relatively large or complex change, feel free to start a discussion by explaining why you chose the solution you did and what alternatives you considered